<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<title>Mali OpenCL SDK v1.1.0: Vectorizing your OpenCL code</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="dynsections.js"></script>
<link href="navtree.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="resize.js"></script>
<script type="text/javascript" src="navtree.js"></script>
<script type="text/javascript">
  $(document).ready(initResizable);
</script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
  $(document).ready(function() { searchBox.OnSelectItem(0); });
</script>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
</head>
<body>
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
 <tbody>
 <tr style="height: 56px;">
  <td id="projectlogo"><img alt="Logo" src="logo.png"/></td>
  <td style="padding-left: 0.5em;">
   <div id="projectname">Mali OpenCL SDK v1.1.0
   </div>
  </td>
 </tr>
 </tbody>
</table>
</div>
<!-- end header part -->
<!-- Generated by Doxygen 1.8.2 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
  <div id="navrow1" class="tabs">
    <ul class="tablist">
      <li><a href="index.html"><span>Home</span></a></li>
      <li class="current"><a href="pages.html"><span>Help&#160;and&#160;Tutorials</span></a></li>
      <li><a href="files.html"><span>Files</span></a></li>
      <li>
        <div id="MSearchBox" class="MSearchBoxInactive">
        <span class="left">
          <img id="MSearchSelect" src="search/mag_sel.png"
               onmouseover="return searchBox.OnSearchSelectShow()"
               onmouseout="return searchBox.OnSearchSelectHide()"
               alt=""/>
          <input type="text" id="MSearchField" value="Search" accesskey="S"
               onfocus="searchBox.OnSearchFieldFocus(true)" 
               onblur="searchBox.OnSearchFieldFocus(false)" 
               onkeyup="searchBox.OnSearchFieldChange(event)"/>
          </span><span class="right">
            <a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
          </span>
        </div>
      </li>
    </ul>
  </div>
</div><!-- top -->
<div id="side-nav" class="ui-resizable side-nav-resizable">
  <div id="nav-tree">
    <div id="nav-tree-contents">
      <div id="nav-sync" class="sync"></div>
    </div>
  </div>
  <div id="splitbar" style="-moz-user-select:none;" 
       class="ui-resizable-handle">
  </div>
</div>
<script type="text/javascript">
$(document).ready(function(){initNavTree('hello_world_vector_tutorial.html','');});
</script>
<div id="doc-content">
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
     onmouseover="return searchBox.OnSearchSelectShow()"
     onmouseout="return searchBox.OnSearchSelectHide()"
     onkeydown="return searchBox.OnSearchSelectKey(event)">
<a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(0)"><span class="SelectionMark">&#160;</span>All</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(1)"><span class="SelectionMark">&#160;</span>Classes</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(2)"><span class="SelectionMark">&#160;</span>Files</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(3)"><span class="SelectionMark">&#160;</span>Functions</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(4)"><span class="SelectionMark">&#160;</span>Variables</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(5)"><span class="SelectionMark">&#160;</span>Macros</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(6)"><span class="SelectionMark">&#160;</span>Pages</a></div>

<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0" 
        name="MSearchResults" id="MSearchResults">
</iframe>
</div>

<div class="header">
  <div class="headertitle">
<div class="title">Vectorizing your OpenCL code </div>  </div>
</div><!--header-->
<div class="contents">
<div class="textblock"><p>How to vectorize the hello world example to increase performance.</p>
<h1><a class="anchor" id="helloWorldVectorTheory"></a>
Vectorization Basics</h1>
<h2><a class="anchor" id="helloWorldVectorPerformance"></a>
Performance Benefits</h2>
<p>If there is hardware support for vectors on your device, multiple arithmetic operations can happen simultaneously. For example, hardware which has 128-bit vector hardware could do arithmetic operations of four 32-bit integers simultaneously.</p>
<p>Often, OpenCL kernels do the same operations on multiple pieces of data. In those cases the code can usually be vectorized such that each kernel does arithmetic calculation on more than one piece of the data, and those arithmetic operations are done simultaneously. Obviously, doing more calculations simultaneously leads to a performance benefit.</p>
<h2><a class="anchor" id="helloWorldVectorQuery"></a>
Querying for Hardware Support</h2>
<p>OpenCL devices can advertise their preferred vector widths for the different OpenCL data types. You can use this information to select a kernel that is optimised for the platform you are running on.</p>
<p>For example, one device may only have hardware support for scalar integers while another could have hardware support for integer vectors of width 4. Two versions of the kernel could be written, one using scalars and one using vectors, with the correct version being selected at runtime.</p>
<p>Here is an example of querying for the preferred integer vector width on a particular device (taken from <a class="el" href="hello__world__vector_8cpp.html">hello_world_vector.cpp</a>):</p>
<div class="fragment"><div class="line">    <span class="comment">/*</span></div>
<div class="line"><span class="comment">     * Query the device to find out it&#39;s prefered integer vector width.</span></div>
<div class="line"><span class="comment">     * Although we are only printing the value here, it can be used to select between</span></div>
<div class="line"><span class="comment">     * different versions of a kernel.</span></div>
<div class="line"><span class="comment">     */</span></div>
<div class="line">    cl_uint integerVectorWidth;</div>
<div class="line">    clGetDeviceInfo(device, CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT, <span class="keyword">sizeof</span>(cl_uint), &amp;integerVectorWidth, NULL);</div>
<div class="line">    cout &lt;&lt; <span class="stringliteral">&quot;Prefered vector width for integers: &quot;</span> &lt;&lt; integerVectorWidth &lt;&lt; endl;</div>
</div><!-- fragment --><p> The same is possible for the other OpenCL data types.</p>
<h2><a class="anchor" id="helloWorldVectorMali"></a>
Mali Hardware Support</h2>
<p>Each Mali-T600 series GPU core has a minimum of two 128-bit wide ALUs (Arithmetic logic units), which are vector capable.</p>
<p>Most operations in the ALUs (e.g. floating point add, floating point multiply, integer addition, integer multiply), can operate on 128-bit vector data (e.g. char16, short8, int4, float4).</p>
<p>Use the querying method above to determine the correct vector size to use for your data type.</p>
<p><b>We recommend the use of vectors wherever possible when using a Mali-T600 series GPU.</b></p>
<h1><a class="anchor" id="helloWorldVectorVectorizing"></a>
Vectorizing The Code</h1>
<p>Other than the following changes in this tutorial, the code for the vector and non-vector code is exactly the same.</p>
<ol type="1">
<li><p class="startli"><b>Modify the kernel to use vectors</b></p>
<p class="startli">Our basic hello world example looked like this (<a class="el" href="hello__world__opencl_8cpp.html">hello_world_opencl.cpp</a>):</p>
<div class="fragment"><div class="line">__kernel <span class="keywordtype">void</span> <a class="code" href="hello__world__opencl_8cl.html#a1e483e8dc82df49f47005c9ede1f0ed7" title="Hello World kernel function.">hello_world_opencl</a>(__global <span class="keywordtype">int</span>* restrict inputA,</div>
<div class="line">                                 __global <span class="keywordtype">int</span>* restrict inputB,</div>
<div class="line">                                 __global <span class="keywordtype">int</span>* restrict output)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">/*</span></div>
<div class="line"><span class="comment">     * Set i to be the ID of the kernel instance.</span></div>
<div class="line"><span class="comment">     * If the global work size (set by clEnqueueNDRangeKernel) is n,</span></div>
<div class="line"><span class="comment">     * then n kernels will be run and i will be in the range [0, n - 1].</span></div>
<div class="line"><span class="comment">     */</span></div>
<div class="line">    <span class="keywordtype">int</span> i = get_global_id(0);</div>
<div class="line"></div>
<div class="line">    <span class="comment">/* Use i as an index into the three arrays. */</span></div>
<div class="line">    output[i] = inputA[i] + inputB[i];</div>
<div class="line">}</div>
</div><!-- fragment --><p> Each instance of this kernel does a single integer addition in one operation. We can vectorise this code to do multiple integer additions in a single operation.</p>
<p class="startli">Because of the vector hardware capabilities of Mali-T600 series GPUs, these vector operations can be exectuted in the same time as a single integer addition.</p>
<p class="startli">We can vectorise like so (<a class="el" href="hello__world__vector_8cl.html">hello_world_vector.cl</a>):</p>
<div class="fragment"><div class="line">__kernel <span class="keywordtype">void</span> <a class="code" href="hello__world__vector_8cl.html#afbd9d391c741f52cb9cd5a53afd57d89" title="Vectorized Hello World kernel function.">hello_world_vector</a>(__global <span class="keywordtype">int</span>* restrict inputA,</div>
<div class="line">                                 __global <span class="keywordtype">int</span>* restrict inputB,</div>
<div class="line">                                 __global <span class="keywordtype">int</span>* restrict output)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">/*</span></div>
<div class="line"><span class="comment">     * We have reduced the global work size (n) by a factor of 4 compared to the hello_world_opencl sample.</span></div>
<div class="line"><span class="comment">     * Therefore, i will now be in the range [0, (n / 4) - 1].</span></div>
<div class="line"><span class="comment">     */</span></div>
<div class="line">    <span class="keywordtype">int</span> i = get_global_id(0);</div>
<div class="line"></div>
<div class="line">    <span class="comment">/*</span></div>
<div class="line"><span class="comment">     * Load 4 integers into &#39;a&#39;.</span></div>
<div class="line"><span class="comment">     * The offset calculation is implicit from the size of the vector load.</span></div>
<div class="line"><span class="comment">     * For vloadN(i, p), the address of the first data loaded would be p + i * N.</span></div>
<div class="line"><span class="comment">     * Load from the data from the address: inputA + i * 4.</span></div>
<div class="line"><span class="comment">     */</span></div>
<div class="line">    int4 a = vload4(i, inputA);</div>
<div class="line">    <span class="comment">/* Do the same for inputB */</span></div>
<div class="line">    int4 b = vload4(i, inputB);</div>
<div class="line"></div>
<div class="line">    <span class="comment">/*</span></div>
<div class="line"><span class="comment">     * Do the vector addition.</span></div>
<div class="line"><span class="comment">     * Store the result at the address: output + i * 4.</span></div>
<div class="line"><span class="comment">     */</span></div>
<div class="line">    vstore4(a + b, i, output);</div>
<div class="line">}</div>
</div><!-- fragment --></li>
<li><p class="startli"><b>Reduce the number of kernel instances</b></p>
<p class="startli">Because each kernel instance is now doing multiple additions, we must reduce the number of kernel instances accordingly (<a class="el" href="hello__world__vector_8cpp.html">hello_world_vector.cpp</a>):</p>
<div class="fragment"><div class="line">    <span class="comment">/*</span></div>
<div class="line"><span class="comment">     * Each instance of our OpenCL kernel now operates on 4 elements of each array so the number of</span></div>
<div class="line"><span class="comment">     * instances needed is the number of elements in the array divided by 4.</span></div>
<div class="line"><span class="comment">     */</span></div>
<div class="line">    <span class="keywordtype">size_t</span> globalWorksize[1] = {arraySize / 4};</div>
<div class="line">    <span class="comment">/* Enqueue the kernel */</span></div>
<div class="line">    <span class="keywordflow">if</span> (!<a class="code" href="common_8cpp.html#a73727ed6f5d821d8c97756404e145644" title="Check an OpenCL error number for errors.">checkSuccess</a>(clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL, globalWorksize, NULL, 0, NULL, &amp;event)))</div>
<div class="line">    {</div>
<div class="line">        <a class="code" href="common_8cpp.html#acd459e74d0cef3c4616c7ce1a5a47f4d" title="Release any OpenCL objects that have been created.">cleanUpOpenCL</a>(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);</div>
<div class="line">        cerr &lt;&lt; <span class="stringliteral">&quot;Failed enqueuing the kernel. &quot;</span> &lt;&lt; __FILE__ &lt;&lt; <span class="stringliteral">&quot;:&quot;</span>&lt;&lt; __LINE__ &lt;&lt; endl;</div>
<div class="line">        <span class="keywordflow">return</span> 1;</div>
<div class="line">    }</div>
</div><!-- fragment --><p> The reduction factor is based on the width of the vectors. For example, if we'd used int8's in the kernel instead of int4's, we'd reduce the global work size by a factor 8.</p>
</li>
</ol>
<h1><a class="anchor" id="helloWorldVectorMoreInformation"></a>
More Information</h1>
<p>For more information have a look at the basic OpenCL code in <a class="el" href="hello__world__opencl_8cpp.html">hello_world_opencl.cpp</a> and <a class="el" href="hello__world__opencl_8cl.html">hello_world_opencl.cl</a> and the vectorized version in <a class="el" href="hello__world__vector_8cpp.html">hello_world_vector.cpp</a> and <a class="el" href="hello__world__vector_8cl.html">hello_world_vector.cl</a>.</p>
<p>Main Tutorial: <a class="el" href="hello_world_tutorial.html">Hello World</a>.</p>
<p>Previous section: <a class="el" href="hello_world_open_c_l_tutorial.html">From C to OpenCL</a>.</p>
<p><a class="el" href="hello_world_running_tutorial.html">Running the Samples</a> </p>
</div></div><!-- contents -->
</div><!-- doc-content -->
<!-- start footer part -->
<div id="nav-path" class="navpath"><!-- id is needed for treeview function! -->
  <ul>
    <li class="navelem"><a class="el" href="tutorials.html">Tutorials</a></li><li class="navelem"><a class="el" href="hello_world_tutorial.html">Hello World</a></li>
    <li class="footer">
        <a href="http://www.arm.com/">(C) ARM Ltd. 2013</a>
    </li>
  </ul>
</div>
</body>
</html>
